Skip to content

ITS: factor out CPU/GPU common code & separate Cell class#15293

Merged
shahor02 merged 4 commits intoAliceO2Group:devfrom
f3sch:its/pr/1026
Apr 24, 2026
Merged

ITS: factor out CPU/GPU common code & separate Cell class#15293
shahor02 merged 4 commits intoAliceO2Group:devfrom
f3sch:its/pr/1026

Conversation

@f3sch
Copy link
Copy Markdown
Collaborator

@f3sch f3sch commented Apr 20, 2026 •
edited
Loading

Based on #15289 (including some commits there).

This reduces the total memory consumed by some percent by splitting the CellSeed class into a CellSeed (with only three cluster ind.) and a TrackSeed (with NLayers cluster ind.):
Figure_1

Additionally, this factors out common functions used by the GPU/CPU code (and actually fixes some GPU stuff) by creating a separate header with commonly used functions for the track/cell fit and the accounting of the final track output.

@f3sch f3sch marked this pull request as ready for review April 20, 2026 12:55
@alibuild
Copy link
Copy Markdown
Collaborator

alibuild commented Apr 20, 2026

Error while checking build/O2/fullCI_slc9 for 5cd39b1 at 2026-04-20 21:09:

## sw/BUILD/O2-latest/log
/sw/SOURCES/O2/slc9_x86-64-slc9_x86-64/0/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h(56): error: calling a __device__ function("o2::track::TrackParametrizationWithError<float> ::resetCovariance(float)") from a __host__ __device__ function("resetTrackCovariance") is not allowed
/sw/SOURCES/O2/slc9_x86-64-slc9_x86-64/0/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h(57): error: calling a __device__ function("_ZNK2o25track29TrackParametrizationWithErrorI1?E6getCovEv") from a __host__ __device__ function("resetTrackCovariance") is not allowed
/sw/SOURCES/O2/slc9_x86-64-slc9_x86-64/0/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h(57): error: calling a __device__ function("_ZN2o25track29TrackParametrizationWithErrorI1?E6setCovE1?1?") from a __host__ __device__ function("resetTrackCovariance") is not allowed
/sw/SOURCES/O2/slc9_x86-64-slc9_x86-64/0/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h(158): error: calling a __device__ function("_ZN2o25track29TrackParametrizationWithErrorI1?E6rotateE1?1?1?") from a __host__ __device__ function("fitTrack") is not allowed
/sw/SOURCES/O2/slc9_x86-64-slc9_x86-64/0/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h(161): error: calling a __device__ function("_ZNK2o24base14PropagatorImplI1?E12propagateToXE1?1?1?1?1?1?1?1?1?") from a __host__ __device__ function("fitTrack") is not allowed
/sw/SOURCES/O2/slc9_x86-64-slc9_x86-64/0/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h(168): error: calling a __device__ function("_ZN2o25track29TrackParametrizationWithErrorI1?E18correctForMaterialE1?1?1?1?") from a __host__ __device__ function("fitTrack") is not allowed
/sw/SOURCES/O2/slc9_x86-64-slc9_x86-64/0/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h(173): error: calling a __device__ function("_ZN2o25track29TrackParametrizationWithErrorI1?E6rotateE1?") from a __host__ __device__ function("fitTrack") is not allowed
/sw/SOURCES/O2/slc9_x86-64-slc9_x86-64/0/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h(176): error: calling a __device__ function("_ZNK2o24base14PropagatorImplI1?E12propagateToXE1?1?1?1?1?1?1?1?") from a __host__ __device__ function("fitTrack") is not allowed
/sw/SOURCES/O2/slc9_x86-64-slc9_x86-64/0/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h(183): error: calling a __device__ function("_ZN2o25track29TrackParametrizationWithErrorI1?E18correctForMaterialE1?1?1?") from a __host__ __device__ function("fitTrack") is not allowed
/sw/SOURCES/O2/slc9_x86-64-slc9_x86-64/0/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h(189): error: calling a __device__ function("_ZNK2o25track29TrackParametrizationWithErrorI1?E21getPredictedChi2QuietE1?1?") from a __host__ __device__ function("fitTrack") is not allowed
/sw/SOURCES/O2/slc9_x86-64-slc9_x86-64/0/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h(194): error: calling a __device__ function("_ZN2o25track29TrackParametrizationWithErrorI1?E6updateE1?1?") from a __host__ __device__ function("fitTrack") is not allowed
/sw/SOURCES/O2/slc9_x86-64-slc9_x86-64/0/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h(198): error: calling a __device__ function("_ZN2o25track20TrackParametrizationI1?E4setYE1?") from a __host__ __device__ function("fitTrack") is not allowed
/sw/SOURCES/O2/slc9_x86-64-slc9_x86-64/0/Detectors/ITSMFT/ITS/tracking/include/ITStracking/TrackHelpers.h(199): error: calling a __device__ function("_ZN2o25track20TrackParametrizationI1?E4setZE1?") from a __host__ __device__ function("fitTrack") is not allowed
ninja: build stopped: subcommand failed.

Full log here.

f3sch added 2 commits April 21, 2026 11:02
@f3sch
ITS: add memory stats
3a30c1f 
Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch>
@f3sch
ITS: clear tracklets after cell finding
f60aefb 
Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch>
@alibuild
Copy link
Copy Markdown
Collaborator

alibuild commented Apr 21, 2026

Error while checking build/O2/fullCI_slc9 for 0adc2f1 at 2026-04-21 11:12:

## sw/BUILD/O2-latest/log
ninja: error: dependency cycle: GPU/GPUTracking/G__O2GPUDataTypes.cxx -> stage/lib64/G__O2TPCFastTransformation_rdict.pcm -> stage/lib64/libO2GPUDataTypes.so -> GPU/GPUTracking/CMakeFiles/O2lib-GPUDataTypes.dir/DataTypes/GPUDataTypesConfig.cxx.o -> cmake_object_order_depends_target_O2lib-GPUDataTypes -> GPU/GPUTracking/G__O2GPUDataTypes.cxx

Full log here.

f3sch added 2 commits April 23, 2026 10:31
@f3sch
ITS: separate into cell and track seed class
000f22e 
Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch>
@f3sch
ITS: factor common functions out
c99b0ad 
Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch>
@f3sch
Copy link
Copy Markdown
Collaborator Author

f3sch commented Apr 24, 2026

ptRatioCN I checked that GPU/CPU remains deterministic and that the CPU output is the same as before. @shahor02 I think this is ok to merge, it should also simplify adding your track projection mode until the cov. matrix stabelizes for the GPU code since it is now shared in TrackHelpers header.

@shahor02 shahor02 merged commit 425f555 into AliceO2Group:dev Apr 24, 2026
12 checks passed
@f3sch f3sch deleted the its/pr/1026 branch April 24, 2026 08:55
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@shahor02 shahor02 shahor02 approved these changes

@fprino fprino Awaiting requested review from fprino fprino is a code owner

@mcoquet642 mcoquet642 Awaiting requested review from mcoquet642 mcoquet642 is a code owner

@mconcas mconcas Awaiting requested review from mconcas mconcas is a code owner

Assignees

No one assigned

Labels

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

3 participants

@f3sch @alibuild @shahor02